#include <cuda_runtime.h>

// CUDA核函数定义
__global__ void addVectors(float *a, float *b, float *c, int n) {
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    if (index < n) {
        c[index] = a[index] + b[index];
    }
}


__global__ void kernelExample() {
    int tid = threadIdx.x;
    __shared__ float sharedData[256];
    // 每个线程写入共享内存
    sharedData[tid] = tid;
    // 同步点：所有线程必须在此等待，直到所有写入完成
    __syncthreads();
    // 所有线程都同步后，可以安全地读取共享内存
    if (tid == 0) {
        // 例如：执行某种基于共享数据的计算
    }
}


__global__ void exampleKernel(int *d_data) {
    __shared__ int sharedData[256];
    int tid = threadIdx.x;

    // 每个线程写入共享内存
    sharedData[tid] = tid;

    // 线程同步，确保所有线程都完成了对共享内存的写入
    __threadfence_block();

    // 现在可以安全地读取共享内存中的数据
    int value = sharedData[tid];

    // 将读取到的值写入全局内存
    d_data[tid] = value;
}


__global__ void kernel1(int *d_data) {
    // 执行一些计算并将结果写入全局内存
    d_data[threadIdx.x] = some_computation();

    // 系统级同步，确保所有之前的全局内存写入操作对所有线程可见
    __threadfence_system();
}

__global__ void kernel2(int *d_data) {
    // 在kernel1的所有全局内存写入操作完成后读取数据
    int value = d_data[threadIdx.x];
    // 执行后续计算
}


// 主函数中启动两个内核
kernel1<<<1, N>>>(d_data);
kernel2<<<1, N>>>(d_data);


__global__ void exampleKernel(int *d_data) {
    // 假设某些条件下需要动态选择同步范围
#ifdef USE_BLOCK_SYNC
    __threadfence_block();
#elif defined(USE_SYSTEM_SYNC)
    __threadfence_system();
#else
    // 注意：以下用法并不推荐，仅用于说明
    __threadfence();
#endif

    // 继续执行后续计算
}


